#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/sort.h>
#include <thrust/execution_policy.h>
#include <thread>
#include <thrust/reduce.h>
#include <thrust/functional.h>
#include <thrust/execution_policy.h>
#include <algorithm>

#define THREAD_PER_BLOCK 1024
__global__ void reduce2(int *d_in, int *d_out){
    __shared__ int sdata[THREAD_PER_BLOCK];
    
    for(int cishu = 0; cishu < 100; cishu++) {
        //each thread loads one element from global memory to shared mem
        unsigned int i = blockIdx.x*blockDim.x+threadIdx.x;
        unsigned int tid = threadIdx.x;
        sdata[tid]=d_in[i];
        __syncthreads();
    
        // do reduction in shared mem
        for(unsigned int s=blockDim.x/2; s>0; s>>=1){
            if(tid < s){
                sdata[tid] = max(sdata[tid], sdata[tid+s]);
            }
            __syncthreads();
        }
        
        // write result for this block to global mem
        if(tid==0)d_out[blockIdx.x]=sdata[tid];
    }
}

__global__ void reduce3(int *d_in, int *d_out){
    __shared__ int sdata[THREAD_PER_BLOCK];
    for(int cishu = 0; cishu < 100; cishu++) {
        //each thread loads one element from global memory to shared mem
        unsigned int i=blockIdx.x*(blockDim.x*2)+threadIdx.x;
        unsigned int tid=threadIdx.x;
        sdata[tid] = max(d_in[i], d_in[i+blockDim.x]);
        __syncthreads();
    
        // do reduction in shared mem
        for(unsigned int s=blockDim.x/2; s>0; s>>=1){
            if(tid < s){
                sdata[tid] = max(sdata[tid], sdata[tid+s]);
            }
            __syncthreads();
        }
        
        // write result for this block to global mem
        if(tid==0)d_out[blockIdx.x]=sdata[tid];
    }
}

int main() {
    int n = 2048;
    int *h_vector_idx = new int[n];
    int *h_vector_score = new int[n];
    int *d_vector_idx = nullptr; cudaMalloc(&d_vector_idx, sizeof(int) * n);
    int *d_vector_score = nullptr; cudaMalloc(&d_vector_score, sizeof(int) * n);

    // 初始化数据
    for (int i = 0; i < n; i++) {
        h_vector_idx[i] = rand();
        h_vector_score[i] = i;
    }
    
    // 将数据从主机端传输到设备端
    // d_data = h_data;

    cudaMemcpy(d_vector_idx, h_vector_idx, sizeof(int) * n, cudaMemcpyHostToDevice);
    // cudaMemcpy(d_vector_score, h_vector_score, sizeof(int) * n, cudaMemcpyHostToDevice);
    clock_t time1,time2;
    time1 = clock();
    reduce3<<<1, THREAD_PER_BLOCK>>>(d_vector_idx, d_vector_score);
    cudaDeviceSynchronize();
    time2 = clock();
    std::cout<<(double)(time2-time1)/CLOCKS_PER_SEC<<std::endl;

    // cudaMemcpy(h_vector_idx, d_vector_idx, sizeof(int) * n, cudaMemcpyDeviceToHost);
    cudaMemcpy(h_vector_score, d_vector_score, sizeof(int) * n, cudaMemcpyDeviceToHost);
    std::sort(h_vector_idx, h_vector_idx + n);

    bool equal = h_vector_score[0] == h_vector_idx[n - 1];
    std::cout << h_vector_score[0] << " " << h_vector_idx[n - 1] << "\n";
    std::cout << equal << "\n";

    return;
}